[Metal] Support assert() #1959

k-ye · 2020-10-15T12:38:13Z

Similar to how we've supported print(), I reserved a dedicated area in the Metal buffer to store assertion failure messages. The assertion message and print messages are placed in the same buffer for simplicity. Here's the buffer memory layout:

// MetalPrintAssertBuffer memory view:
//
// +------------------------+ \
// | AssertRecorderData     | |
// +------------------------+ | -> for assert(), kMetalAssertBufferSize
// | PrintMsg for assert()  | |
// +------------------------+ /
// | PrintMsgAllocator      | \
// +------------------------+ |
// |                        | |
// |                        | |
// | a queue of PrintMsgs   | | -> for print()
// | for print()            | |
// | ... ...                | |
// |                        | |
// +------------------------+ /

We only store the error message of the first assertion failure. This is done by using AssertRecorder::mark_first_failure(). Under the hood it's just an atomic flag.

Sample generated code:

@ti.kernel
def assert_formatted():
    for i in x:
        assert x[i] == 0, 'x[%d] expect=%d got=%d' % (i, 0, x[i])

if (!(tmp22)) {
    if (assert_rec_.mark_first_failure()) {
      assert_rec_.set_num_args(4);
      PrintMsg tmp23_msg_(assert_rec_.msg_buf_addr(), 4);
      tmp23_msg_.pm_set_str(/*i=*/0, 2);  // 2 is the index in the string table. We don't store the actual string on GPU.
      tmp23_msg_.pm_set_i32(1, tmp2);
      tmp23_msg_.pm_set_i32(2, tmp11);
      tmp23_msg_.pm_set_i32(3, tmp20);
    }
    return;
  }

To answer @archibate question in #1955 (review), hopefully it's now clear? These backends can all share format_error_message(), which does the formatting of %.

This PR is not urgent at all. Please take your time and consider it low priority if you have other things to deal with...

Related issue = #1955

[Click here for the format server]

codecov · 2020-10-15T13:19:47Z

Codecov Report

Merging #1959 into master will increase coverage by 0.94%.
The diff coverage is n/a.

@@            Coverage Diff             @@
##           master    #1959      +/-   ##
==========================================
+ Coverage   42.69%   43.64%   +0.94%     
==========================================
  Files          45       45              
  Lines        6436     6225     -211     
  Branches     1106     1106              
==========================================
- Hits         2748     2717      -31     
+ Misses       3513     3334     -179     
+ Partials      175      174       -1

Impacted Files	Coverage Δ
python/taichi/lang/ast_checker.py	`70.58% <0.00%> (-1.64%)`	⬇️
python/taichi/testing.py	`75.00% <0.00%> (-0.72%)`	⬇️
python/taichi/lang/linalg.py	`89.33% <0.00%> (-0.67%)`	⬇️
python/taichi/lang/meta.py	`62.82% <0.00%> (-0.48%)`	⬇️
python/taichi/lang/__init__.py	`40.45% <0.00%> (-0.33%)`	⬇️
python/taichi/misc/util.py	`18.27% <0.00%> (-0.24%)`	⬇️
python/taichi/misc/task.py	`0.00% <0.00%> (ø)`
python/taichi/tools/patterns.py	`0.00% <0.00%> (ø)`
python/taichi/lang/kernel.py	`71.39% <0.00%> (+0.13%)`	⬆️
python/taichi/misc/gui.py	`8.89% <0.00%> (+0.18%)`	⬆️
... and 9 more

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update 706c519...c9172cf. Read the comment docs.

yuanming-hu

Great! LGTM. Hopefully we can handle assertions uniformly on all backends.

I guess for printing, on CPU and CUDA the built-in synchronous printing support is still helpful. On CUDA the printf support is kind of limited, so we may want to use the print buffer as well in the future. We may also want to support Python-style (or C++20 std::format-style) formatting on the backends :-)

[Metal] Support assert()

75e5607

k-ye requested review from yuanming-hu, archibate and taichi-gardener October 15, 2020 12:38

simplify

c9172cf

yuanming-hu approved these changes Oct 16, 2020

View reviewed changes

k-ye merged commit 6fea4c2 into taichi-dev:master Oct 16, 2020

k-ye deleted the mtl-assert branch October 16, 2020 11:35

yuanming-hu mentioned this pull request Oct 17, 2020

[release] v0.7.0 #1971

Merged

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[Metal] Support assert() #1959

[Metal] Support assert() #1959

k-ye commented Oct 15, 2020 •

edited

Loading

codecov bot commented Oct 15, 2020 •

edited

Loading

yuanming-hu left a comment

[Metal] Support assert() #1959

[Metal] Support assert() #1959

Conversation

k-ye commented Oct 15, 2020 • edited Loading

codecov bot commented Oct 15, 2020 • edited Loading

Codecov Report

yuanming-hu left a comment

Choose a reason for hiding this comment

k-ye commented Oct 15, 2020 •

edited

Loading

codecov bot commented Oct 15, 2020 •

edited

Loading